#include "..\..\include\pcvf.h"
#include <CL\cl.h>

extern cl_context			g_device_context;
extern cl_command_queue		g_device_command_queue;
extern cl_kernel			g_kernel_list[pcvClsMAX][pcvFuncMax];

const char* pcvMean_8u_C1Rs_Kernel_Src1[] = {
"__kernel void pcvMean_8u_C1Rs_Kernel(",
"	__global unsigned char*	img_src,",
"	__global unsigned char*	img_dst,",
"	int						size_width,",
"	int						roi_width)",
"{",
"	int				y = get_global_id(0);",
"	int				x;",
"	int				offset = y * size_width;",
"	int				sum;",
"	for (x = 0; x < roi_width; ++x)",
"	{",
"		sum = img_src[offset + x];",
"		sum += img_src[offset + 1 + x];",
"		sum += img_src[offset + 2 + x];",
"		sum += img_src[offset + size_width + x];",
"		sum += img_src[offset + size_width + 1 + x];",
"		sum += img_src[offset + size_width + 2 + x];",
"		sum += img_src[offset + size_width + size_width + x];",
"		sum += img_src[offset + size_width + size_width + 1 + x];",
"		sum += img_src[offset + size_width + size_width + 2 + x];",
"		img_dst[offset + size_width + 1 + x] = sum / 9;",
"	}",
"}"
};

const char* pcvMean_8u_C1Rs_Kernel_Src[] = {
"__kernel void pcvMean_8u_C1Rs_Kernel(",
"	__global unsigned char*	img_src,",
"	__global unsigned char*	img_dst,",
"	int						size_width,",
"	int						roi_width)",
"{",
"	int				y = get_global_id(0);",
"	int				x = 0;",
"	int				offset = y * size_width;",
"	int				offset_0 = offset + 2;",
"	int				offset_1 = offset + size_width + 2;",
"	int				offset_2 = offset + size_width + size_width + 2;",
"	int				sum_buffer[3];",
"	sum_buffer[1] = img_src[offset];",
"	sum_buffer[1] += img_src[offset + size_width];",
"	sum_buffer[1] += img_src[offset + size_width + size_width];",
"	sum_buffer[2] = img_src[offset + 1];",
"	sum_buffer[2] += img_src[offset + size_width + 1];",
"	sum_buffer[2] += img_src[offset + size_width + size_width + 1];",
"	for (x = 0; x < roi_width; ++x)",
"	{",
"		sum_buffer[x % 3] = img_src[offset_0 + x];",
"		sum_buffer[x % 3] += img_src[offset_1 + x];",
"		sum_buffer[x % 3] += img_src[offset_2 + x];",
"		img_dst[offset + size_width + 1 + x] = (sum_buffer[0] + sum_buffer[1] + sum_buffer[2]) / 9;",
"	}",
"}"
};
//"	img_dst[offset] = add_sat(img_src[offset], img_src1[offset]);",

PCVAPI(PcvStatus, pcvMean_8u_C1Rs_Build, (
	))
{
	cl_int					errcode = CL_SUCCESS;

	if (g_device_context == NULL ||
		g_kernel_list[pcvClsF][pcvFuncFMean_8u_C1Rs] != NULL)
	{
		return pcvStsErr;
	}

	// OpenCL supports runtime code compilation. First we build the OpenCL
	// program from source code
	cl_program				program = clCreateProgramWithSource(g_device_context, 27, pcvMean_8u_C1Rs_Kernel_Src, NULL, &errcode);
	if (errcode != CL_SUCCESS)
	{
		return pcvStsErr;
	}

	errcode = clBuildProgram(program, 0, NULL, NULL, NULL, NULL);
	if (errcode != CL_SUCCESS)
	{
		clReleaseProgram(program);
		return pcvStsErr;
	}

	// Create a handle to the compiled OpenCL function (Kernel)
	g_kernel_list[pcvClsF][pcvFuncFMean_8u_C1Rs] = clCreateKernel(program, "pcvMean_8u_C1Rs_Kernel", &errcode);
	if (errcode != CL_SUCCESS)
	{
		clReleaseProgram(program);
		return pcvStsErr;
	}

	errcode = clReleaseProgram(program);
	if (errcode != CL_SUCCESS)
	{
		return pcvStsErr;
	}

	return pcvStsNoErr;
}

PCVAPI(PcvStatus, pcvMean_8u_C1Rs, (
	const void*				device_img_src,
	void*					device_img_dst,
	const PcvSize			host_size,
	const PcvRect			host_roi
	))
{
	cl_int					errcode = CL_SUCCESS;

	if (g_device_command_queue == NULL)
	{
		return pcvStsNoInitErr;
	}

	cl_kernel				kernel = g_kernel_list[pcvClsF][pcvFuncFMean_8u_C1Rs];
	if (kernel == NULL)
	{
		return pcvStsNotBuildErr;
	}

	// In the next step we associate the GPU memory with the Kernel arguments
	int						offset = (host_roi.y - 1) * host_size.width + host_roi.x - 1;
	unsigned char*			img_src_offset = (unsigned char*)device_img_src + offset;
	errcode = clSetKernelArg(kernel, 0, sizeof(cl_mem), (void*)&img_src_offset);
	if (errcode != CL_SUCCESS)
	{
		return pcvStsErr;
	}

	unsigned char*			img_dst_offset = (unsigned char*)device_img_dst + offset;
	errcode = clSetKernelArg(kernel, 1, sizeof(cl_mem), (void*)&img_dst_offset);
	if (errcode != CL_SUCCESS)
	{
		return pcvStsErr;
	}

	errcode = clSetKernelArg(kernel, 2, sizeof(int), (void*)&(host_size.width));
	if (errcode != CL_SUCCESS)
	{
		return pcvStsErr;
	}

	errcode = clSetKernelArg(kernel, 3, sizeof(int), (void*)&(host_roi.width));
	if (errcode != CL_SUCCESS)
	{
		return pcvStsErr;
	}

	// Then we launch the Kernel on the GPU
	size_t					work_size[1] = {host_roi.height - 2};
	cl_event				event_ret;
	errcode = clEnqueueNDRangeKernel(g_device_command_queue, kernel, 1, NULL, work_size, NULL, 0, NULL, &event_ret);
	if (errcode != CL_SUCCESS)
	{
		return pcvStsErr;
	}

	// Wait for Kernel execution complete
	errcode = clWaitForEvents(1, &event_ret);
	if (errcode != CL_SUCCESS)
	{
		return pcvStsErr;
	}

	cl_event_info			event_info = CL_EVENT_COMMAND_EXECUTION_STATUS;
	cl_int					event_status;
	errcode = clGetEventInfo(event_ret, event_info, sizeof(cl_int), &event_status, NULL);
	if (errcode != CL_SUCCESS || event_status != CL_COMPLETE)
	{
		return pcvStsErr;
	}

	return pcvStsNoErr;
}
